#include "OclLapack.h"
#include "../BindKernelArgs.h"
#include <algorithm>
#include "../ClMem.h"

using namespace std;

const char* OclLapack::programScaleSrc = 
"__kernel void scale1(float sa, __global float* x, int offsetX) { \n"
"	// TODO use vectorization \n"
"	int i = get_global_id(0) + offsetX; \n"
"	x[i] *= sa; \n"
"} \n"
"\n"
"__kernel void scaleN(float sa, __global float* x, int offsetX, int incX) { \n"
"	int i = get_global_id(0) * incX + offsetX; \n"
"	x[i] *= sa; \n"
"} \n"
"\n";

void OclLapack::sscal(int n, float sa, cl_ptr x, int incX) {
	Measure measure("sscal-ocl");
	
	cl_kernel kernel;
	if (incX == 1) {
		kernel = scale1;
		BindKernelArgs(scale1) << sa << x.handle << x.offset;
	} else {
		if (incX < 0) {
			x.offset += (1 - n) * incX;
		}
		kernel = scaleN;
		BindKernelArgs(scaleN) << sa << x.handle << x.offset << incX;
	}
	
	size_t count = n;
	cl_event event;
    CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &count, NULL, 0, NULL, &event));
	ClEvent(event).wait();
}

int OclLapack::sscal(int *n, float *sa, float *sx, int *incx) {
	Measure measure("sscal");
	
	size_t sizeX = 1 + (*n - 1) * abs(*incx);
	ClMem<float> memX(queue, sizeX, sx);
	
	sscal(*n, *sa, cl_ptr(memX), *incx);
	
	memX.beginRead(sx, 0, sizeX).wait();
	
	return 0;
}
